#include <iostream>
#include <cuda_runtime.h>
#include "../common/common.h"
#define DEVICENUM 0
#define BDIMX 32 // 方形共享内存行数
#define BDIMY 32 // 方形共享内存列数
#define IPAD 1

__global__ void setRowReadRow(int *out)
{
    __shared__ int tile[BDIMY][BDIMX];                         // 设置方形静态共享内存
    unsigned int idx = threadIdx.y * blockDim.x + threadIdx.x; // 由于只有一个block，直接计算线程编号即可

    tile[threadIdx.y][threadIdx.x] = idx; // 将线程编号当成值赋值给共享内存，按照行主序设置，无冲突
    __syncthreads();                      // 等待所有线程完成

    out[idx] = tile[threadIdx.y][threadIdx.x]; // 共享内存按照行主序读取到gpu数组中，无冲突
}

__global__ void setColReadCol(int *out)
{
    __shared__ int tile[BDIMY][BDIMX];                         // 设置方形静态共享内存
    unsigned int idx = threadIdx.y * blockDim.x + threadIdx.x; // 计算线程编号即可

    tile[threadIdx.x][threadIdx.y] = idx; // 将线程编号当成值赋值给共享内存，按照列主序设置，BDIMY个冲突，需串行访问存储体
    __syncthreads();                      // 等待所有线程完成

    out[idx] = tile[threadIdx.x][threadIdx.y]; // 共享内存按列主序读取到gpu数组中，BDIMY个冲突，需串行访问存储体
}

__global__ void setRowReadCol(int *out)
{
    __shared__ int tile[BDIMY][BDIMX];                         // 设置方形静态共享内存
    unsigned int idx = threadIdx.y * blockDim.x + threadIdx.x; // 计算线程编号即可

    tile[threadIdx.y][threadIdx.x] = idx; // 将线程编号当成值赋值给共享内存，按照行主序设置
    __syncthreads();                      // 等待所有线程完成

    out[idx] = tile[threadIdx.x][threadIdx.y]; // 共享内存按列主序读取到gpu数组中，BDIMY个冲突，需串行访问存储体
}

__global__ void setRowReadColPad(int *out)
{
    __shared__ int tile[BDIMY][BDIMX + IPAD];                  // 设置方形静态共享内存,此时多了一列
    unsigned int idx = threadIdx.y * blockDim.x + threadIdx.x; // 计算线程编号

    tile[threadIdx.y][threadIdx.x] = idx;      // 将线程编号当成值赋值给共享内存，按照行主序设置
    __syncthreads();                           // 等待所有线程完成
    out[idx] = tile[threadIdx.x][threadIdx.y]; // 虽然按列主序读取到gpu数组中，但由于进行了填充，此时也是无冲突访问
}

int main(int argc, char **argv)
{
    cudaSetDevice(DEVICENUM);
    int nElem = BDIMX * BDIMY;
    size_t nBytes = nElem * sizeof(int);

    // 设置用于接收方形共享内存的gpu数组
    int *device_array_res = nullptr;
    CHECK(cudaMalloc((int **)&device_array_res, nBytes));
    CHECK(cudaMemset(device_array_res, 0, nBytes));

    // 设置用于接收gpu结果数组的cpu指针
    int *gpu_res_to_cpu = (int *)malloc(nBytes);
    memset(gpu_res_to_cpu, 0, nBytes);

    dim3 gridSize(1, 1);
    dim3 blockSize(BDIMX, BDIMY);

    // 1. 方形共享内存按照行设置，按照行读取
    setRowReadRow<<<gridSize, blockSize>>>(device_array_res);
    CHECK(cudaMemcpy(gpu_res_to_cpu, device_array_res, nBytes, cudaMemcpyDeviceToHost));

    // 2. 方形共享内存按照列设置，按照列读取
    setColReadCol<<<gridSize, blockSize>>>(device_array_res);
    CHECK(cudaMemcpy(gpu_res_to_cpu, device_array_res, nBytes, cudaMemcpyDeviceToHost));

    // 3. 方形共享内存按照行设置，按照列读取
    setRowReadCol<<<gridSize, blockSize>>>(device_array_res);
    CHECK(cudaMemcpy(gpu_res_to_cpu, device_array_res, nBytes, cudaMemcpyDeviceToHost));

    // 4. 使用填充解决存储体冲突
    setRowReadColPad<<<gridSize, blockSize>>>(device_array_res);
    CHECK(cudaMemcpy(gpu_res_to_cpu, device_array_res, nBytes, cudaMemcpyDeviceToHost));

    cudaFree(device_array_res);
    free(gpu_res_to_cpu);

    return 0;
}